// TILE_X holds the number of blocks per tile in x direction
// TILE_Y holds the number of blocks per tile in the y direction

#define fast_norm3( v ) (v).w = native_rsqrt((v).x*(v).x + (v).y*(v).y + (v).z+(v).z); (v).x *= (v).w; (v).y *= (v).w; (v).z *= (v).w; (poot).w = 0.0f;
#define fast_norm3_t( temp, v ) (temp) = native_rsqrt((v).x*(v).x + (v).y*(v).y + (v).z*(v).z ); (v).x *= (temp); (v).y *= (temp); (v).z *= (temp);

#define fast_dot3( v1, v2 ) ((v1).x*(v2).x + (v1).y*(v2).y + (v1).z*(v2).z)
#define fast_sub3( r, v1, v2 ) (r).x = (v1).x - (v2).x; (r).y = (v1).y - (v2).y; (r).z = (v1).z - (v2).z;
#define fast_add3( r, v1, v2 ) (r).x = (v1).x + (v2).x; (r).y = (v1).y + (v2).y; (r).z = (v1).z + (v2).z;
#define fast_min3( r, v1, v2 ) (r).x = fmin((v1).x, (v2).x); (r).y = fmin((v1).y, (v2).y); (r).z = fmin((v1).z, (v2).z);
#define fast_max3( r, v1, v2 ) (r).x = fmax((v1).x, (v2).x); (r).y = fmax((v1).y, (v2).y); (r).z = fmax((v1).z, (v2).z);

#if THREADS_PER_TILE < 32
#error you aren't doing this right!
#endif

typedef struct {
	float x;
	float y;
	float z;
} float3h;

/*int boxIntersect( float3h min1, float3h max1, float3h min2, float3h max2 ) {
	return ( min1.x < max2.x )
	&& ( min1.y < max2.y )
	&& ( min1.z < max2.z )
	&& ( max1.x > min2.x )
	&& ( max1.y > min2.y )
	&& ( max1.z > min2.z );
}*/

int lightCulled(float4 light, 
				float3h minCoords,
				float3h maxCoords,
				float zNear,
				float Q1, 
				float Q2,
				float C,
				float D,
				float boxSize) {

	if ( light.w < 0.99f ) {
		// it's a point light, cannot be culled
		return 0;
	}
	
	/* project a bounding box of the light into clip space */
	if ( light.z - boxSize > -zNear ) {
		// if the backside of light's bounding box
		// is in front of the zNear plane
		// then the light should be clipped
		return 1;
	}
	
	float3h lightMin, lightMax;
	float div_front = native_recip( light.z + boxSize );
	float div_back	= native_recip( light.z - boxSize );
	
	if ( light.z + boxSize > -zNear ) {
		// solves the problem of when front plane of bounding box
		// is behind zNear
		div_front = native_recip(-zNear);
	}
	
	// MIN X
	lightMin.x = -Q1 * ( light.x - boxSize );
	lightMin.x *= lightMin.x > 0.0f ? div_front : div_back;
	if ( lightMin.x > maxCoords.x ) return 1;

	// MIN Y
	lightMin.y = -Q2 * ( light.y - boxSize );
	lightMin.y *= lightMin.y > 0.0f ? div_front : div_back;
	if ( lightMin.y > maxCoords.y ) return 1;

	// MAX X
	lightMax.x = -Q1 * ( light.x + boxSize );
	lightMax.x *= lightMax.x > 0.0f ? div_back : div_front;
	if ( lightMax.x < minCoords.x ) return 1;

	// MAX Y
	lightMax.y = -Q2 * ( light.y + boxSize );
	lightMax.y *= lightMax.y > 0.0f ? div_back : div_front;
	if ( lightMax.y < minCoords.y ) return 1;

	// MIN Z
	lightMin.z = -C - D * div_front;
	if ( lightMin.z > maxCoords.z ) return 1;

	// MAX Z
	lightMax.z = -C - D * div_back;
	if ( lightMax.z < minCoords.z ) return 1;
	
	return 0;
											
}

#define PREFIX(i) \
if ( threadID - i >= 0 ) { \
	sum += prefix[ threadID - i ]; \
	barrier(CLK_LOCAL_MEM_FENCE); \
	prefix[ threadID ] = sum; \
	barrier(CLK_LOCAL_MEM_FENCE);\
}\

int prefix_sum(int needsStore, __local int *prefix, int threadID) {

	int sum = needsStore;
	prefix[ threadID ] = sum;
	barrier(CLK_LOCAL_MEM_FENCE);
	
#if THREADS_PER_TILE > 1
	PREFIX(1);
#endif
#if THREADS_PER_TILE > 2
	PREFIX(2);
#endif
#if THREADS_PER_TILE > 4
	PREFIX(4);
#endif
#if THREADS_PER_TILE > 8
	PREFIX(8);
#endif
#if THREADS_PER_TILE > 16
	PREFIX(16);
#endif
#if THREADS_PER_TILE > 32
	PREFIX(32);
#endif
#if THREADS_PER_TILE > 64
	PREFIX(64);
#endif
#if THREADS_PER_TILE > 128
	PREFIX(128);
#endif
	return prefix[ THREADS_PER_TILE - 1];	
}


int prefix_sum_trivial(int needsStore, __local int *prefix, int threadID) {

	prefix[threadID] = needsStore;
	barrier(CLK_LOCAL_MEM_FENCE);
	if ( threadID == 0 ) {
		int accum = 0;
		for (int i=0; i<THREADS_PER_TILE; i++) {
			accum += prefix[i];
			prefix[i] = accum;
		}
	}
	barrier(CLK_LOCAL_MEM_FENCE);
	int num = prefix[THREADS_PER_TILE-1];
	return num;

}

__kernel void binLights(	__read_only	 image2d_t input_image,
					const int2 numTiles,
					const int2 numBlocks,
					__global float4 *lights,
					const int numLights,
					__global int *lightBins,
					__global int *binCounts,
					const		float		zNear,
					const		float		Q1, // 2 * zNear / ( right - left )
					const		float		Q2, // 2 * zNear / ( top - bottom )
					const		float		C,  // - (zFar + zNear) / (zFar - zNear)
					const		float		D,  // - (2 zFar zNear) / (zFar - zNear)
					sampler_t	nearestSampler,
					int2 imageSize) {

	// step 1.
	
	__local int		localPrefix[THREADS_PER_TILE];
	__local int		localIndices[THREADS_PER_TILE];
	__local	int		localReorderBuffer[THREADS_PER_TILE];
	
	int3 gid = { get_global_id(0), get_global_id(1), get_global_id(2) };
	
	float3h minCoords, maxCoords;
	
	int min_block_x = gid.y * (TILE_X);
	int max_block_x	= min_block_x + (TILE_X);
	int min_block_y = gid.z * (TILE_Y);
	int max_block_y	= min_block_y + (TILE_Y);
	max_block_x = min( max_block_x, numBlocks.x - 1);
	max_block_y = min( max_block_y, numBlocks.y - 1);

	int threadID	= get_local_id(0);
	int binID		= ( gid.y + get_global_id(2) * numTiles.x );

	__local float foundMin;
	__local float foundMax;

	if ( threadID == 0 ) {
				
		float locateMin = 1.0f;
		float locateMax = -1.0f;
				
		for (int i=min_block_x; i < max_block_x; i++) {
			for (int j=min_block_y; j < max_block_y; j++) {
				int2 coords = { i, j };
				float4 values = read_imagef( input_image, nearestSampler, coords );
				locateMin = fmin( locateMin, values.x );
				locateMax = fmax( locateMax, values.y );
			}
		}
		
		foundMin = locateMin;
		foundMax = locateMax;
				
	}
	
	barrier(CLK_LOCAL_MEM_FENCE);
	minCoords.z = foundMin;
	maxCoords.z = foundMax;
	
	if ( minCoords.z >= 0.9999f ) {
		if ( threadID == 0 ) {
			binCounts[binID] = 0;
		}
		return;
	}
	
	// normalized device coordinates
	minCoords.x = 2.0f * ((float)(min_block_x	* BLOCK_X)  / (float)(imageSize.x)) - 1.0f;
	minCoords.y = 2.0f * ((float)(min_block_y	* BLOCK_Y)  / (float)(imageSize.y)) - 1.0f;
	maxCoords.x = 2.0f * ((float)(max_block_x 	* BLOCK_X)  / (float)(imageSize.x)) - 1.0f;
	maxCoords.y = 2.0f * ((float)(max_block_y 	* BLOCK_Y)  / (float)(imageSize.y)) - 1.0f;
	// the z coordinates were written in ndc from the last compute shader


	// step 2, bin the lights
	int binOffset = numLights * binID;

	int storeCount = 0;	
	int buffCount = 0;
	
	for (int i=threadID; i<numLights+threadID; i+=THREADS_PER_TILE) {
		
		float4 currentLight;
		int needsStore = 0;
		
		if ( i < numLights ) {
			currentLight = lights[i];
			needsStore = !lightCulled(currentLight, minCoords, maxCoords, zNear, Q1, Q2, C, D, 0.2f);
		}
						
		// compact the lights that need storing
		int numberToStore = prefix_sum(needsStore, localPrefix, threadID);
		
		// can we append to buffer without draining?
		if ( numberToStore + buffCount >= THREADS_PER_TILE ) {
			// drain buffer
			if ( threadID < buffCount ) {
				lightBins[binOffset + storeCount + threadID] = localReorderBuffer[threadID];
			}
			storeCount += buffCount;
			buffCount = 0;	
		}

		if ( needsStore ) {
			int storeIndex = localPrefix[threadID] - 1;
			localIndices[storeIndex] = i;
		}

		barrier(CLK_LOCAL_MEM_FENCE);		
		
		if ( threadID < numberToStore ) {
			int index = localIndices[threadID];
			localReorderBuffer[buffCount + threadID] = index;
		}
				
		// increment the store count by how many lights we just stored
		buffCount += numberToStore;
		
	}
		
	if ( threadID < buffCount ) {
		// drain buffer
		lightBins[binOffset + storeCount + threadID] = localReorderBuffer[threadID];
		storeCount += buffCount;
		buffCount = 0;
	}

	if ( threadID == 0 ) {
		binCounts[binID] = storeCount;
	}
	
}